NVGPUOps
nvgpu.cluster_arrive (triton::nvgpu::ClusterArriveOp)
语法:
operation ::= `nvgpu.cluster_arrive` attr-dict
属性:
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
| relaxed | ::mlir::IntegerAttr | 1-bit 无符号整数属性 |
nvgpu.cluster_id (triton::nvgpu::ClusterCTAIdOp)
语法:
operation ::= `nvgpu.cluster_id` attr-dict
特征:AlwaysSpeculatableImplTrait
接口:ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)
效果:MemoryEffects::Effect{}
结果:
| 结果 | 描述 |
|---|---|
| result | 32-bit 无符号整数 |
nvgpu.cluster_wait (triton::nvgpu::ClusterWaitOp)
语法:
operation ::= `nvgpu.cluster_wait` attr-dict
nvgpu.fence_async_shared (triton::nvgpu::FenceAsyncSharedOp)
语法:
operation ::= `nvgpu.fence_async_shared` attr-dict
属性:
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
| bCluster | ::mlir::BoolAttr | 布尔属性 |
nvgpu.stmatrix (triton::nvgpu::StoreMatrixOp)
语法:
operation ::= `nvgpu.stmatrix` operands attr-dict `:` type(operands)
接口:MemoryEffectOpInterface (MemoryEffectOpInterface)
效果:MemoryEffects::Effect{MemoryEffects::Write on ::mlir::SideEffects::DefaultResource}
操作:
| 操作 | 描述 |
|---|---|
| addr | 地址空间 3 中的 LLVM 指针 |
| datas | 32-bit 无符号整数的可变参数 |
nvgpu.wgmma_commit_group (triton::nvgpu::WGMMACommitGroupOp)
语法:
operation ::= `nvgpu.wgmma_commit_group` attr-dict
nvgpu.wgmma_fence (triton::nvgpu::WGMMAFenceOp)
语法:
operation ::= `nvgpu.wgmma_fence` attr-dict
nvgpu.wgmma (triton::nvgpu::WGMMAOp)
语法:
operation ::= `nvgpu.wgmma` $opA `,` $opB (`,` $opC^)? attr-dict `:` functional-type(operands, $res)
Attributes: 属性:
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
| m | ::mlir::IntegerAttr | 32-bit 无符号整数属性 |
| n | ::mlir::IntegerAttr | 32-bit 无符号整数属性 |
| k | ::mlir::IntegerAttr | 32-bit 无符号整数属性 |
| eltTypeC | ::mlir::triton::nvgpu::WGMMAEltTypeAttr | wgmma 操作数类型为 's8', 's32', 'e4m3', 'e5m2', 'f16', 'bf16', 'tf32', 或者 'f32' |
| eltTypeA | ::mlir::triton::nvgpu::WGMMAEltTypeAttr | wgmma 操作数类型为 's8', 's32', 'e4m3', 'e5m2', 'f16', 'bf16', 'tf32', 或者 'f32' |
| eltTypeB | ::mlir::triton::nvgpu::WGMMAEltTypeAttr | wgmma 操作数类型为 's8', 's32', 'e4m3', 'e5m2', 'f16', 'bf16', 'tf32', 或者 'f32' |
| layoutA | ::mlir::triton::nvgpu::WGMMALayoutAttr | wgmma 布局,可以为 'row'(行)或 'col'(列) |
| layoutB | ::mlir::triton::nvgpu::WGMMALayoutAttr | wgmma 布局,可以为 'row'(行)或 'col'(列) |
操作:
| 操作 | 描述 |
|---|---|
| opA | wgmma operand A/B type wgmma 操作 A/B 类型 |
| opB | wgmma operand A/B type wgmma 操作 A/B 类型 |
| opC | LLVM structure type LLVM 结构体类型 |
结果:
| 结果 | 描述 |
|---|---|
| res | LLVM 结构体类型 |
nvgpu.wgmma_wait_group (triton::nvgpu::WGMMAWaitGroupOp)
语法:
operation ::= `nvgpu.wgmma_wait_group` $input attr-dict `:` type($input)
接口:InferTypeOpInterface
属性:
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
| pendings | ::mlir::IntegerAttr | 32-bit 无符号整数属性 |
操作:
| 操作 | 描述 |
|---|---|
| input | LLVM structure type LLVM 结构体类型 |
结果:
| 结果 | 描述 |
|---|---|
| output | LLVM structure type LLVM 结构体类型 |